#include <cutil_inline.h>
#include "config.cuh"
#include <io.h>

#define MODE 1

char *DATAFILE;
int N;
int CSC;

void parseData(string rec, int* closeSetCount, int i, ha_t* ha)
{
	char * cstr, *p;
	cstr = new char[rec.size()+1];
	strcpy(cstr, rec.c_str());

	float pos = 0, duration = 0, speed = 0, time = 0;
	int baseId = 0, no = 0, etype = 0;

#if MODE ==0
	p=strtok (cstr,"\t");
	no = atoi(p); //no
	p=strtok(NULL,"\t");
	time = (float)atof(p);//time
	p=strtok(NULL,"\t");
	baseId = atoi(p)-1; //baseid
	pos = baseId*2+1;//pos
	p=strtok(NULL,"\t");
	duration = (float)atof(p);//duration
	p=strtok(NULL,"\t");
	speed = (float)atof(p);//speed
	etype=INIT;
#else
	p=strtok (cstr,"\t");
	no = atoi(p); //no
	p=strtok(NULL,"\t");
	time = (float)atof(p);//time
	p=strtok(NULL,"\t");
	duration = (float)atof(p);//duration
	p=strtok(NULL,"\t");
	pos = (float)atof(p);//pos
	baseId = (int)(pos/2);//baseid
	p=strtok(NULL,"\t");
	speed = (float)atof(p);//speed
	etype=INIT;
#endif
	int endBaseId;
	if(pos+duration*speed/3600>40)
		endBaseId=19;
	else
		endBaseId=(pos+duration*speed/3600)/2;

	ha->init[i]=*closeSetCount;
	ha->baseId[*closeSetCount]=baseId;
	ha->duration[*closeSetCount]=duration;
	ha->etype[*closeSetCount]=etype;
	ha->no[*closeSetCount]=no;
	ha->pos[*closeSetCount]=pos;
	ha->speed[*closeSetCount]=speed;
	ha->time[*closeSetCount]=time;
	*closeSetCount+=((endBaseId-baseId)*2+2);
}

int init(ha_t* ha, da_t* da)
{
	ifstream fin;
	fin.open(DATAFILE);
	string rec;

	ha->baseId = (int*)malloc(CSC*sizeof(int));
	ha->duration = (float*)malloc(CSC*sizeof(float));
	ha->speed = (float*)malloc(CSC*sizeof(float));
	ha->pos = (float*)malloc(CSC*sizeof(float));
	ha->no = (int*)malloc(CSC*sizeof(int));
	ha->time = (float*)malloc(CSC*sizeof(float));
	ha->etype = (int*)malloc(CSC*sizeof(int));
	ha->init = (int*)malloc(N*sizeof(int));

	int i = 0;
	int counter = 0;
	while(!fin.eof()){
		getline(fin,rec);
		parseData(rec, &counter, i, ha);
		i++;
	}
	fin.close();

	cudaMalloc((void**)&(da->baseId), counter*sizeof(int));
	cudaMalloc((void**)&(da->index),counter*sizeof(int));
	cudaMalloc((void**)&(da->time),counter*sizeof(float));
	cudaMalloc((void**)&(da->no), counter*sizeof(int));
	cudaMalloc((void**)&(da->etype), counter*sizeof(int));
	cudaMalloc((void**)&(da->sum), counter*sizeof(int));
	cudaMalloc((void**)&(da->duration), counter*sizeof(float));
	cudaMalloc((void**)&(da->speed), counter*sizeof(float));
	cudaMalloc((void**)&(da->pos), counter*sizeof(float));
	cudaMalloc((void**)&(da->init), N*sizeof(int));
	cudaMalloc((void**)&(da->detect), sizeof(int));

	cudaMemcpy(da->time, ha->time, counter*sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(da->baseId,ha->baseId,counter*sizeof(int),cudaMemcpyHostToDevice);
	cudaMemcpy(da->duration,ha->duration,counter*sizeof(float),cudaMemcpyHostToDevice);
	cudaMemcpy(da->speed,ha->speed,counter*sizeof(float),cudaMemcpyHostToDevice);
	cudaMemcpy(da->pos,ha->pos,counter*sizeof(float),cudaMemcpyHostToDevice);
	cudaMemcpy(da->no,ha->no,counter*sizeof(int),cudaMemcpyHostToDevice);
	cudaMemcpy(da->etype,ha->etype,counter*sizeof(int),cudaMemcpyHostToDevice);
	cudaMemcpy(da->init,ha->init, N*sizeof(int), cudaMemcpyHostToDevice);

	free(ha->time);
	free(ha->baseId);
	free(ha->duration);
	free(ha->speed);
	free(ha->pos);
	free(ha->no);
	free(ha->etype);
	free(ha->init);

	return counter;
}

/*memcpy in launch kernel and main produce different effects possibly because
the extern indentifier*/
extern "C" void launch_init_kernel(da_t da, int count);
extern "C" void launch_sort_kernel(da_t da, int count);
extern "C" void launch_scan_kernel(da_t da, int count);
extern "C" bool launch_corr_kernel(da_t da, int count);

int main(int argc, char** argv)
{
	DATAFILE = argv[1];
	N = atoi(argv[2]);
	CSC = 10*N;

	ha_t ha;
	da_t da;

	CUevent start, stop;
	CUevent start_o, stop_o; //overall

	remove("corr2.txt");
	remove("fini1.txt");
	remove("fini2.txt");

	timer_start(start, stop);
	//INIT_HOST
	int counter = init(&ha, &da);
	int n = counter;
	cout<<counter<<" "<<sizeof(ha)<<endl;

	//INIT_kernel after this process, d_speed, d_pos, d_duration, d_init is freed
	launch_init_kernel(da, N);
	
	//SORT_KERNEL after this process, d_time, d_index is freed
	launch_sort_kernel(da, counter);
	
	//SCAN_KERNEL	
	launch_scan_kernel(da, counter);

	//CORRECTION_KERNEL
#if VERSION == 0
	bool detect = launch_corr_kernel(da, counter);
#else
	bool detect = launch_corr_kernel(da, N);
#endif

	//REPEATLY SCAN and CORRECT
	cout<<detect<<" ";
	int aa = 0;
	while(detect == true){
		launch_scan_kernel(da, counter);
#if VERSION == 0
		detect = launch_corr_kernel(da, counter);
#else
		detect = launch_corr_kernel(da, N);
#endif
		aa++;
	}
	timer_stop(start, stop);
//	cout<<aa<<" ";
	//printinfo("fini2.txt");
	//printinfofini("fini1.txt", counter);
	int *hno = (int*)malloc(counter*sizeof(int));
	int res = 0;
	cudaMemcpy(hno, da.no, counter*sizeof(int), cudaMemcpyDeviceToHost);
	for(int i=0; i<counter; i++)
		if(hno[i]<0)
			res++;
	cout<<res<<endl;
}